/* Template project which demonstrates the basics on how to setup a project 
* example application.
* Host code.
*/

// includes, system
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>

// includes, project
#include <cutil.h>

// includes, kernels
#include <template_kernel.cu>

#include "QP_Utility_header.cu"
#include "QP_Utility.cu"
#include "template_gold.cu"
#include "scan.cu"



unsigned int timer=0;
void startTime()
{
	timer=0;
	CUT_SAFE_CALL( cutCreateTimer( &timer));
    CUT_SAFE_CALL( cutStartTimer( timer));
}
void endTime(char *op)
{
	CUT_SAFE_CALL( cutStopTimer( timer));
	printf( "%s processing time, %f ,ms\n", op, cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutDeleteTimer( timer));
}

void endMULTime(char *op, float times)
{
	CUT_SAFE_CALL( cutStopTimer( timer));
	printf( "%s processing time, %f ,ms\n", op, cutGetTimerValue( timer)/times);
    CUT_SAFE_CALL( cutDeleteTimer( timer));
}


void initilizeTwoArray(int2 *Sorted, int rLen, int2 *unSorted, int uLen)
{
	generateSort(Sorted, TEST_MAX, rLen, 0);
	generateRand(unSorted,TEST_MAX, uLen, 1);
}

void initilizeOneArray(int2 *R, int rLen, int numUpdate)
{
	psgenUpdate(R, rLen, numUpdate, TEST_MAX, 0);
}



void cpu_sortTwoArray(int2 *Sorted, int rLen, int2 *unSorted, int uLen, int2* output)
{
	char info[256];
	sprintf(info, "cpusort, %d, %d",rLen,uLen);
	startTime();
	qsort(unSorted,uLen,sizeof(int2),compareValue);
	endTime("sort the unsorted region");	
	startTime();
	cpu_mergeSort(Sorted, rLen, unSorted, uLen, output);
	endTime(info);	
}
void cpu_ResortTwoArray(int2 *Sorted, int rLen, int2 *unSorted, int uLen, int2* output)
{
	char info[256];
	sprintf(info, "cpuResort, %d, %d",rLen,uLen);
	startTime();
	cpu_reSort(Sorted, rLen, unSorted, uLen, output);
	endTime(info);	
}

void gpu_sortTwoArray(int2 *Sorted, int rLen, int2 *unSorted, int uLen, int2* output)
{
	char info[256];
	sprintf(info, "gpusort, %d, %d",rLen,uLen);
	startTime();
	cpu_mergeSort(Sorted, rLen, unSorted, uLen, output);
	endTime(info);	
}

void testTwoArray()
{	
	int rLen=1024*2*1024;
	int uLen=256*1024;
	int2 *Sorted=(int2*)malloc(sizeof(int2)*rLen);
	int2 *unSorted=(int2*)malloc(sizeof(int2)*uLen);
	int totalLen=rLen+uLen;
	int2 *output=(int2*)malloc(sizeof(int2)*totalLen);
	initilizeTwoArray(Sorted, rLen, unSorted, uLen);
	cpu_sortTwoArray(Sorted, rLen, unSorted, uLen,output);
	//cpu_ResortTwoArray(Sorted, rLen, unSorted, uLen,output);
	//gpu_sortTwoArray(Sorted, rLen, unSorted, uLen,output);
	checkSorted(output, totalLen);
	free(output);
	free(unSorted);
	free(Sorted);
}

void cpu_sortOneArray(int2 *R, int rLen, int2* output, char* info)
{
	startTime();
	//cpu_sort(R, rLen, output);
	adaptiveSort(R,rLen,output);
	endTime(info);
}

void testOneArray(int numUpdateRatio)
{	
	int rLen=1024;
	int numUpdate=rLen*numUpdateRatio/100;
	int2 *R=(int2*)malloc(sizeof(int2)*rLen);
	int2 *output=(int2*)malloc(sizeof(int2)*rLen);
	char info[256];
	sprintf(info, "cpusort, %d, U, %d",rLen, numUpdate);
	initilizeOneArray(R, rLen, numUpdate);
	cpu_sortOneArray(R,rLen,output, info);
	checkSorted(output, rLen);
	free(output);
	free(R);
}

void gpu_NagaSort(int rLen,int randomNum)
{
	printf("rLen, %d, random number, %d\n", rLen, randomNum);
	int2 *h_R=(int2 *)malloc(sizeof(int2)*rLen);
	int i=0;
	for(i=0;i<rLen;i++)
	{
		h_R[i].y=i;
		h_R[i].x=i;
	}
	randomize(h_R, rLen, randomNum);
	int numBlock=rLen/NUM_THREAD_PER_BLOCK;
	int2 *d_R;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_R, rLen*sizeof(int2) )); 
	CUDA_SAFE_CALL( cudaMemcpy( d_R, h_R,	rLen*sizeof(int2), cudaMemcpyHostToDevice) );
	gpuPrintInt2(d_R,rLen,"input");
	int2 *d_output;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_output, rLen*sizeof(int2) )); 
	int *d_pmin;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_pmin, numBlock*sizeof(int) )); 
	int *d_I;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_I, rLen*sizeof(int) )); 
	int *d_A;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_A, rLen*sizeof(int) )); 
	int *d_B;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_B, rLen*sizeof(int) )); 
	int *d_C;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_C, rLen*sizeof(int) )); 
	int *d_one;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_one, (rLen)*sizeof(int) ));
	//CUDA_SAFE_CALL( cudaMemset(  d_one, 0, (rLen+1)*sizeof(int) )); //the last slot for sum.
	int *d_zero;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_zero, (rLen)*sizeof(int) ));
	int *d_oneSum;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_oneSum, (rLen)*sizeof(int) ));
	int *d_zeroSum;
	CUDA_SAFE_CALL( cudaMalloc( (void**) & d_zeroSum, (rLen)*sizeof(int) ));
	int *h_pmin=(int *)malloc(sizeof(int)*numBlock);
	int* h_oneSize=(int*)malloc(sizeof(int));
	
	int2* oldDR=d_R;
	int old_rLen=rLen;
	int THRESHOLD=32*NUM_THREAD_PER_BLOCK;
	preallocBlockSums(rLen);
	int pass=0;
	startTime();
	while(rLen>THRESHOLD)
	{
		//step1, prefix min from right ot left.
		numBlock=ceil((double)rLen/(double)NUM_THREAD_PER_BLOCK);
		
		//startTime();
		dim3  ps_Thread( NUM_THREAD_PER_BLOCK, 1, 1);
		dim3  ps_grid( numBlock, 1 , 1);
		right_min_scan_partialMin<<< ps_grid, ps_Thread>>>(d_R, d_I,rLen, d_pmin);	
		CUDA_SAFE_CALL( cudaMemcpy( h_pmin, d_pmin, numBlock*sizeof(int) , cudaMemcpyDeviceToHost) );
		int tempValue=TEST_MAX;
		int curValue=0;
		for(int i=numBlock-1;i>=0;i--)
		{
			curValue=h_pmin[i];
			h_pmin[i]=tempValue;
			tempValue=MIN(tempValue,curValue);
		}
		CUDA_SAFE_CALL( cudaMemcpy( d_pmin, h_pmin,	numBlock*sizeof(int), cudaMemcpyHostToDevice) );
		gpuPrint(d_pmin, numBlock, "d_output");
		right_min_scan_fullMin<<< ps_grid, ps_Thread>>>(d_I, d_pmin, rLen, d_B);
		gpuPrint(d_I, rLen, "d_I");
		gpuPrint(d_B, rLen, "d_B");

		//compute prefix min from left to right.
		left_min_scan_partialMin<<< ps_grid, ps_Thread>>>(d_I,d_B,rLen, d_pmin);	
		CUDA_SAFE_CALL( cudaMemcpy( h_pmin, d_pmin, numBlock*sizeof(int) , cudaMemcpyDeviceToHost) );
		gpuPrint(d_pmin, numBlock, "d_output1");
		tempValue=TEST_MAX;
		curValue=0;
		for(int i=0;i<numBlock;i++)
		{
			curValue=h_pmin[i];
			h_pmin[i]=tempValue;
			tempValue=MIN(tempValue,curValue);
		}
		CUDA_SAFE_CALL( cudaMemcpy( d_pmin, h_pmin,	numBlock*sizeof(int), cudaMemcpyHostToDevice) );
		gpuPrint(d_pmin, numBlock, "d_output");
		left_min_scan_fullMin<<< ps_grid, ps_Thread>>>(d_I, d_pmin, d_B, rLen, d_one, d_zero);
		gpuPrint(d_one, rLen, "d_one");
		gpuPrint(d_zero, rLen, "d_zero");
		//endTime("scans");

		//prefix sum on the C (non-zero elements, zeroed elements).
		//startTime();
		prescanArray(d_oneSum, d_one,old_rLen);
		prescanArray(d_zeroSum, d_zero,old_rLen);
		//endTime("one prefix sum");
		gpuPrint(d_oneSum, rLen, "d_oneSum");
		gpuPrint(d_zeroSum, rLen, "d_zeroSum");
		
		CUDA_SAFE_CALL( cudaMemcpy( h_oneSize, d_oneSum+(rLen-1), sizeof(int) , cudaMemcpyDeviceToHost) );
		printf("pass, %d, h_oneSize, %d, unsorted, %d\n",pass, *h_oneSize, rLen-*h_oneSize);

		aScatter<<< ps_grid, ps_Thread>>>(d_R, d_output, d_one, d_oneSum, d_zeroSum, *h_oneSize, rLen);
		//gpuPrintInt2(d_output,rLen,"d_output after scatter");	
		CUDA_SAFE_CALL( cudaMemcpy( d_R, d_output,	rLen*sizeof(int2), cudaMemcpyDeviceToDevice) );
		gpuPrintInt2(d_R,old_rLen,"sorted");		
		//update the size information.
		rLen=rLen-*h_oneSize;
		d_R=d_R+*h_oneSize;
		//
		pass++;
	}
	endTime("NagaSort");
	/*int2 *h_output=(int2*)malloc(sizeof(int2*)*old_rLen);
	CUDA_SAFE_CALL( cudaMemcpy( h_output, oldDR, sizeof(int2)*old_rLen, cudaMemcpyDeviceToHost) );
	checkSorted(h_output,old_rLen);
	deallocBlockSums();
	
	free(h_output);*/
	CUDA_SAFE_CALL( cudaFree( d_one ));
	CUDA_SAFE_CALL( cudaFree( d_zero ));
	CUDA_SAFE_CALL( cudaFree( d_I ));
	CUDA_SAFE_CALL( cudaFree( d_A ));
	CUDA_SAFE_CALL( cudaFree( d_B ));
	CUDA_SAFE_CALL( cudaFree( d_C ));	
}

void test_cpu_nagaSort(int rLen, int randomNum)
{
	//rLen=6;
	//int a[6]={5,   6 ,      8 ,      7  ,     9  ,     10};
	int2 *R=(int2 *)malloc(sizeof(int2)*rLen);
	int2 *output=(int2 *)malloc(sizeof(int2)*rLen);
	int i=0;
	for(i=0;i<rLen;i++)
	{
		R[i].y=i;
		R[i].x=i;
	}
	randomize(R, rLen, randomNum);
	int pass=0;
	printf("rLen, %d, random number, %d\n", rLen, randomNum);
	int *I=(int *)malloc(sizeof(int)*rLen);
	int *A=(int *)malloc(sizeof(int)*rLen);
	int *B=(int *)malloc(sizeof(int)*rLen);
	int *C=(int *)malloc(sizeof(int)*rLen);
	startTime();
	cpu_NagaSort(R,rLen,output, pass, I, A, B, C);
	endTime("NagaSort");
	checkSorted(R,rLen);

	/*for(i=0;i<rLen;i++)
	{
		R[i].y=i;
		R[i].x=i;
	}
	randomize(R, rLen, randomNum);
	startTime();
	qsort(R,rLen,sizeof(int2),compareValue);
	endTime("QuickSort");*/

	free(C);
	free(B);
	free(A);
	free(I);


}

/*
int main( int argc, char** argv) 
{
	//testTwoArray();
	//for(int r=1;r<=64;r=r*2)
	//int r=1;
	//testOneArray(r);
	//for(int r=1;r<=64;r=r*2)
	int r=1;
	int rLen=1024;
	if(argc==2) r=atoi(argv[1]);
	if(argc==3) 
	{
		rLen=atoi(argv[1])*1024;
		r=atoi(argv[2]);
	}
	gpu_NagaSort(rLen,r);
	//for(int r=1;r<=64;r=r*2)
	test_cpu_nagaSort(rLen, r);

	

    
    return EXIT_SUCCESS;
}

*/